#ifndef _TEMPLATE_KERNEL_H_
#define _TEMPLATE_KERNEL_H_
#include <QP_Utility_header.cu>
#include <cutil.h>
#include <stdio.h>
#include <stdlib.h>


__global__ void 
right_min_scan_partialMin(int2* d_input, int* d_I, int d_len, int* d_output)
{
	__shared__ int pMin[NUM_THREAD_PER_BLOCK];
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int bid=bx+by*gridDim.x;
	int tid=tx+ty*blockDim.x;
	int gid=bid*NUM_THREAD_PER_BLOCK+tid;
	//load the data into the shared memory
	if(gid<d_len)
	{
		int2 tempInt2=d_input[gid];
		pMin[tid]=tempInt2.y;
		d_I[gid]=tempInt2.y;
	}
	else
	{
		pMin[tid]=TEST_MAX;
	}	
	__syncthreads();
    //sum
	for (int d = NUM_THREAD_PER_BLOCK>>1; d > 0; d >>= 1)
    {
        __syncthreads();

        if (tid < d)      
        {
            int ai = tid+d;
            int bi = tid;
			if(pMin[bi]>pMin[ai])
            pMin[bi] = pMin[ai];//write the min of the two.
        }
    }
	__syncthreads();
	if(tid==0)
		d_output[bid]=pMin[0];
}

__global__ void 
right_min_scan_fullMin(int* d_I, int* d_partialMin, int d_len, int* d_B)
{
	__shared__ int pMin[NUM_THREAD_PER_BLOCK];
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int bid=bx+by*gridDim.x;
	int tid=tx+ty*blockDim.x;
	int gid=bid*NUM_THREAD_PER_BLOCK+tid;
	//load the data into the shared memory
	int iValue;
	if(gid<d_len)
	{
		iValue=d_I[gid];
		pMin[tid]=iValue;
	}
	else
		pMin[tid]=TEST_MAX;
	int previousValue=d_partialMin[bid];
	__syncthreads();
    //min
	if(tid==0)
	{
		int curValue=0;
		for(int i=NUM_THREAD_PER_BLOCK-1;i>=0;i--)
		{
			curValue=pMin[i];
			pMin[i]=previousValue;
			previousValue=MIN(curValue,previousValue);
		}
	}
	__syncthreads();
	
	if(gid<d_len)
	{
		if(iValue<=pMin[tid])
			previousValue=1;
		else
			previousValue=0;
		d_B[gid]=previousValue;
	}
}


__global__ void 
left_min_scan_partialMin(int* d_I, int* d_B, int d_len, int* d_output)
{
	__shared__ int pMin[NUM_THREAD_PER_BLOCK];
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int bid=bx+by*gridDim.x;
	int tid=tx+ty*blockDim.x;
	int gid=bid*NUM_THREAD_PER_BLOCK+tid;
	//load the data into the shared memory
	int iValue;
	int bValue;
	if(gid<d_len)
	{
		iValue=d_I[gid];
		bValue=d_B[gid];
		if(bValue==0)
			pMin[tid]=iValue;
		else
			pMin[tid]=TEST_MAX;
	}
	else
		pMin[tid]=TEST_MAX;
	__syncthreads();
    //min
	for (int d = NUM_THREAD_PER_BLOCK>>1; d > 0; d >>= 1)
    {
        __syncthreads();

        if (tid < d)      
        {
            int ai = tid+d;
            int bi = tid;
			if(pMin[bi]>pMin[ai])
            pMin[bi] = pMin[ai];//write the min of the two.
        }
    }
	__syncthreads();
	if(tid==0)
		d_output[bid]=pMin[0];
}


__global__ void 
left_min_scan_fullMin(int* d_I, int* d_partialMin, int* d_B, int d_len, int* d_one, int *d_zero)
{
	__shared__ int pMin[NUM_THREAD_PER_BLOCK];
	__shared__ int pB[NUM_THREAD_PER_BLOCK];
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int bid=bx+by*gridDim.x;
	int tid=tx+ty*blockDim.x;
	int gid=bid*NUM_THREAD_PER_BLOCK+tid;
	//load the data into the shared memory
	int iVlaue;
	if(gid<d_len)
	{
		iVlaue=d_I[gid];
		pMin[tid]=iVlaue;
		pB[tid]=d_B[gid];
	}
	else
		pMin[tid]=TEST_MAX;
	int previousValue=d_partialMin[bid];
	__syncthreads();
    //min
	if(tid==0)
	{
		int curValue=0;
		for(int i=0;i<NUM_THREAD_PER_BLOCK;i++)
		{
			if(pB[i]==0)
				curValue=pMin[i];
			else
				curValue=TEST_MAX;
			pMin[i]=previousValue;
			previousValue=MIN(curValue,previousValue);
		}
	}
	__syncthreads();
	
	if(gid<d_len)
	{	
		if(iVlaue<=pMin[tid] && pB[tid]==1)
			previousValue=1;
		else
			previousValue=0;
		d_one[gid]=previousValue;
		d_zero[gid]=1-previousValue;
	}
}
__global__
void aScatter(int2* d_R, int2* d_output, int* d_one, int* d_oneSum, int* d_zeroSum, int oneSize,int d_len)
{
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int bid=bx+by*gridDim.x;
	int tid=tx+ty*blockDim.x;
	int gid=bid*NUM_THREAD_PER_BLOCK+tid;
	if(gid<d_len)
	{
		int bValue=d_one[gid];
		int oneValue=d_oneSum[gid];
		int zeroValue=d_zeroSum[gid];
		int dest=0;
		if(bValue==1)
			dest=oneValue;
		else
		{
			dest=zeroValue;
			dest+=oneSize;
		}
		d_output[dest]=d_R[gid];		
	}
}





#endif // #ifndef _TEMPLATE_KERNEL_H_
